/*
 *  Copyright 2010 INFN - APE group
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

#pragma once

#include <cuqu/config.h>

namespace cuqu
{
    namespace detail
    {
        struct atomic {
            unsigned int m_val;
        };

        typedef struct atomic atomic_t;

#define CUQU_INIT_ATOMIC { 0 }

        static __inline__ __device__ void atomic_set(atomic_t *ptr, unsigned int val)
        {
            // return value discarded
            atomicExch(&ptr->m_val, val);
        }

        // return old value
        static __inline__ __device__ unsigned int atomic_mod_inc(atomic_t *ptr, unsigned int modulus)
        {
            return atomicInc(&ptr->m_val, modulus);
        }

        // return old value
        static __inline__ __device__ unsigned int atomic_inc(atomic_t *ptr)
        {
            return atomicInc(&ptr->m_val, 0xffffffffU);
        }

        static __inline__ __device__ unsigned int atomic_add(atomic_t *ptr, unsigned int n)
        {
            return atomicAdd(&ptr->m_val, n);
        }

        static __inline__ __device__ unsigned int atomic_read(atomic_t *ptr)
        {
            unsigned int ret;
            ret = *(volatile unsigned int*)&(ptr->m_val);
            return ret;
        }
    } // end namespace detail

} // end namespace cuqu

/*
 * Local variables:
 *  mode: c++
 *  c-indent-level: 4
 *  c-basic-offset: 4
 *  tab-width: 4
 *  indent-tabs-mode: nil
 * End:
 */
